Generated by Cython 3.0.2
Yellow lines hint at Python interaction.
Click on a line that starts with a "+" to see the C code that Cython generated for it.
Raw output: _le_convolution.c
+001: # cython: infer_types=True, wraparound=False, nonecheck=False, boundscheck=False, cdivision=True, language_level=3, profile=False, autogen_pxd=False
__pyx_t_4 = __Pyx_PyDict_NewPresized(0); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 1, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); if (PyDict_SetItem(__pyx_d, __pyx_n_s_test, __pyx_t_4) < 0) __PYX_ERR(0, 1, __pyx_L1_error) __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
002:
+003: import numpy as np
__pyx_t_7 = __Pyx_ImportDottedModule(__pyx_n_s_numpy, NULL); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 3, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_np, __pyx_t_7) < 0) __PYX_ERR(0, 3, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0;
+004: import time
__pyx_t_7 = __Pyx_ImportDottedModule(__pyx_n_s_time, NULL); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 4, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_time, __pyx_t_7) < 0) __PYX_ERR(0, 4, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0;
005:
006:
007: cimport numpy as np
008:
009: from cython.parallel import parallel, prange
010:
011: from libc.math cimport cos, sin
012:
+013: from .__interpolation_tools__ import check_image, value2array
__pyx_t_7 = PyList_New(2); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __Pyx_INCREF(__pyx_n_s_check_image); __Pyx_GIVEREF(__pyx_n_s_check_image); if (__Pyx_PyList_SET_ITEM(__pyx_t_7, 0, __pyx_n_s_check_image)) __PYX_ERR(0, 13, __pyx_L1_error); __Pyx_INCREF(__pyx_n_s_value2array); __Pyx_GIVEREF(__pyx_n_s_value2array); if (__Pyx_PyList_SET_ITEM(__pyx_t_7, 1, __pyx_n_s_value2array)) __PYX_ERR(0, 13, __pyx_L1_error); __pyx_t_4 = __Pyx_Import(__pyx_n_s_interpolation_tools, __pyx_t_7, 1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __pyx_t_7 = __Pyx_ImportFrom(__pyx_t_4, __pyx_n_s_check_image); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_check_image, __pyx_t_7) < 0) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __pyx_t_7 = __Pyx_ImportFrom(__pyx_t_4, __pyx_n_s_value2array); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_value2array, __pyx_t_7) < 0) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
+014: from .__liquid_engine__ import LiquidEngine
__pyx_t_4 = PyList_New(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 14, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_INCREF(__pyx_n_s_LiquidEngine); __Pyx_GIVEREF(__pyx_n_s_LiquidEngine); if (__Pyx_PyList_SET_ITEM(__pyx_t_4, 0, __pyx_n_s_LiquidEngine)) __PYX_ERR(0, 14, __pyx_L1_error); __pyx_t_7 = __Pyx_Import(__pyx_n_s_liquid_engine, __pyx_t_4, 1); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 14, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_ImportFrom(__pyx_t_7, __pyx_n_s_LiquidEngine); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 14, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); if (PyDict_SetItem(__pyx_d, __pyx_n_s_LiquidEngine, __pyx_t_4) < 0) __PYX_ERR(0, 14, __pyx_L1_error) __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0;
+015: from .__opencl__ import cl, cl_array
__pyx_t_7 = PyList_New(2); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 15, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __Pyx_INCREF(__pyx_n_s_cl); __Pyx_GIVEREF(__pyx_n_s_cl); if (__Pyx_PyList_SET_ITEM(__pyx_t_7, 0, __pyx_n_s_cl)) __PYX_ERR(0, 15, __pyx_L1_error); __Pyx_INCREF(__pyx_n_s_cl_array); __Pyx_GIVEREF(__pyx_n_s_cl_array); if (__Pyx_PyList_SET_ITEM(__pyx_t_7, 1, __pyx_n_s_cl_array)) __PYX_ERR(0, 15, __pyx_L1_error); __pyx_t_4 = __Pyx_Import(__pyx_n_s_opencl_2, __pyx_t_7, 1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 15, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __pyx_t_7 = __Pyx_ImportFrom(__pyx_t_4, __pyx_n_s_cl); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 15, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_cl, __pyx_t_7) < 0) __PYX_ERR(0, 15, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __pyx_t_7 = __Pyx_ImportFrom(__pyx_t_4, __pyx_n_s_cl_array); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 15, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_cl_array, __pyx_t_7) < 0) __PYX_ERR(0, 15, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
016:
017:
+018: class Convolution(LiquidEngine):
__Pyx_GetModuleGlobalName(__pyx_t_4, __pyx_n_s_LiquidEngine); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 18, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_7 = PyTuple_New(1); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 18, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __Pyx_GIVEREF(__pyx_t_4); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_7, 0, __pyx_t_4)) __PYX_ERR(0, 18, __pyx_L1_error); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PEP560_update_bases(__pyx_t_7); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 18, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_5 = __Pyx_CalculateMetaclass(NULL, __pyx_t_4); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 18, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __pyx_t_9 = __Pyx_Py3MetaclassPrepare(__pyx_t_5, __pyx_t_4, __pyx_n_s_Convolution, __pyx_n_s_Convolution, (PyObject *) NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_kp_s_2D_convolution); if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 18, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); if (__pyx_t_4 != __pyx_t_7) { if (unlikely((PyDict_SetItemString(__pyx_t_9, "__orig_bases__", __pyx_t_7) < 0))) __PYX_ERR(0, 18, __pyx_L1_error) } __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __pyx_t_7 = PyList_New(0); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 18, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); /* … */ __pyx_t_10 = __Pyx_Py3ClassCreate(__pyx_t_5, __pyx_n_s_Convolution, __pyx_t_4, __pyx_t_9, NULL, 0, 0); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 18, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_10); if (__Pyx_CyFunction_InitClassCell(__pyx_t_7, __pyx_t_10) < 0) __PYX_ERR(0, 18, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; if (PyDict_SetItem(__pyx_d, __pyx_n_s_Convolution, __pyx_t_10) < 0) __PYX_ERR(0, 18, __pyx_L1_error) __Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0; __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; __Pyx_DECREF(__pyx_t_5); __pyx_t_5 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
019: """
020: 2D convolution
021: """
022:
+023: def __init__(self, clear_benchmarks=False, testing=False):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_1__init__(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_1__init__ = {"__init__", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_1__init__, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_1__init__(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
PyObject *__pyx_v_self = 0;
PyObject *__pyx_v_clear_benchmarks = 0;
PyObject *__pyx_v_testing = 0;
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("__init__ (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 23, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_clear_benchmarks,&__pyx_n_s_testing,0};
PyObject* values[3] = {0,0,0};
values[1] = __Pyx_Arg_NewRef_FASTCALL(((PyObject *)((PyObject *)Py_False)));
values[2] = __Pyx_Arg_NewRef_FASTCALL(((PyObject *)((PyObject *)Py_False)));
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 23, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (kw_args > 0) {
PyObject* value = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_clear_benchmarks);
if (value) { values[1] = __Pyx_Arg_NewRef_FASTCALL(value); kw_args--; }
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 23, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (kw_args > 0) {
PyObject* value = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_testing);
if (value) { values[2] = __Pyx_Arg_NewRef_FASTCALL(value); kw_args--; }
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 23, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "__init__") < 0)) __PYX_ERR(0, 23, __pyx_L3_error)
}
} else {
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
break;
default: goto __pyx_L5_argtuple_error;
}
}
__pyx_v_self = values[0];
__pyx_v_clear_benchmarks = values[1];
__pyx_v_testing = values[2];
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("__init__", 0, 1, 3, __pyx_nargs); __PYX_ERR(0, 23, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution.__init__", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution___init__(__pyx_self, __pyx_v_self, __pyx_v_clear_benchmarks, __pyx_v_testing);
/* function exit code */
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution___init__(CYTHON_UNUSED PyObject *__pyx_self, PyObject *__pyx_v_self, PyObject *__pyx_v_clear_benchmarks, PyObject *__pyx_v_testing) {
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("__init__", 0);
/* … */
/* function exit code */
__pyx_r = Py_None; __Pyx_INCREF(Py_None);
goto __pyx_L0;
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution.__init__", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__23 = PyTuple_Pack(3, __pyx_n_s_self, __pyx_n_s_clear_benchmarks, __pyx_n_s_testing); if (unlikely(!__pyx_tuple__23)) __PYX_ERR(0, 23, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__23);
__Pyx_GIVEREF(__pyx_tuple__23);
__pyx_codeobj__24 = (PyObject*)__Pyx_PyCode_New(3, 0, 0, 3, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__23, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_convoluti, __pyx_n_s_init, 23, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__24)) __PYX_ERR(0, 23, __pyx_L1_error)
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_1__init__, 0, __pyx_n_s_Convolution___init, NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_d, ((PyObject *)__pyx_codeobj__24)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 23, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
__Pyx_INCREF(__pyx_t_10);
PyList_Append(__pyx_t_7, __pyx_t_10);
__Pyx_GIVEREF(__pyx_t_10);
__Pyx_CyFunction_SetDefaultsTuple(__pyx_t_10, __pyx_tuple__25);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_init, __pyx_t_10) < 0) __PYX_ERR(0, 23, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
__pyx_tuple__25 = PyTuple_Pack(2, ((PyObject *)Py_False), ((PyObject *)Py_False)); if (unlikely(!__pyx_tuple__25)) __PYX_ERR(0, 23, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__25);
__Pyx_GIVEREF(__pyx_tuple__25);
+024: self._designation = "Conv2D"
if (__Pyx_PyObject_SetAttrStr(__pyx_v_self, __pyx_n_s_designation, __pyx_n_u_Conv2D) < 0) __PYX_ERR(0, 24, __pyx_L1_error)
+025: super().__init__(clear_benchmarks=clear_benchmarks, testing=testing,
__pyx_t_1 = __Pyx_CyFunction_GetClassObj(__pyx_self); if (!__pyx_t_1) { PyErr_SetString(PyExc_SystemError, "super(): empty __class__ cell"); __PYX_ERR(0, 25, __pyx_L1_error) } __Pyx_INCREF(__pyx_t_1); __pyx_t_2 = PyTuple_New(2); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 25, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_2, 0, __pyx_t_1)) __PYX_ERR(0, 25, __pyx_L1_error); __Pyx_INCREF(__pyx_v_self); __Pyx_GIVEREF(__pyx_v_self); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_2, 1, __pyx_v_self)) __PYX_ERR(0, 25, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyObject_Call(__pyx_builtin_super, __pyx_t_2, NULL); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 25, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_init); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 25, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(8); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 25, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_clear_benchmarks, __pyx_v_clear_benchmarks) < 0) __PYX_ERR(0, 25, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_testing, __pyx_v_testing) < 0) __PYX_ERR(0, 25, __pyx_L1_error) /* … */ __pyx_t_3 = __Pyx_PyObject_Call(__pyx_t_2, __pyx_empty_tuple, __pyx_t_1); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 25, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0;
+026: opencl_=True, unthreaded_=True, threaded_=True, threaded_static_=True,
if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_opencl, Py_True) < 0) __PYX_ERR(0, 25, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_unthreaded, Py_True) < 0) __PYX_ERR(0, 25, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_threaded, Py_True) < 0) __PYX_ERR(0, 25, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_threaded_static, Py_True) < 0) __PYX_ERR(0, 25, __pyx_L1_error)
+027: threaded_dynamic_=True, threaded_guided_=True)
if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_threaded_dynamic, Py_True) < 0) __PYX_ERR(0, 25, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_threaded_guided, Py_True) < 0) __PYX_ERR(0, 25, __pyx_L1_error)
+028: self._default_benchmarks = {'OpenCL': {"(['shape(100, 100)', 'shape(5, 5)'], {})": [250000, 0.009790166019229218, 0.003298707975773141, 0.003652999992482364], "(['shape(500, 500)', 'shape(11, 11)'], {})": [30250000, 0.007485667010769248, 0.007351625012233853, 0.007932541979243979]}, 'Threaded': {"(['shape(100, 100)', 'shape(5, 5)'], {})": [250000, 0.00022562500089406967, 0.00018245799583382905, 0.00017095799557864666], "(['shape(500, 500)', 'shape(11, 11)'], {})": [30250000, 0.007719500019447878, 0.0060990000201854855, 0.008289082994451746]}, 'Threaded_dynamic': {"(['shape(100, 100)', 'shape(5, 5)'], {})": [250000, 0.0001590839819982648, 0.0001595419889781624, 0.00020799998310394585], "(['shape(500, 500)', 'shape(11, 11)'], {})": [30250000, 0.005388707999372855, 0.005681082984665409, 0.005188624985748902]}, 'Threaded_guided': {"(['shape(100, 100)', 'shape(5, 5)'], {})": [250000, 0.00018945799092762172, 0.0001723750028759241, 0.00017620899598114192], "(['shape(500, 500)', 'shape(11, 11)'], {})": [30250000, 0.005136792024131864, 0.005925459001446143, 0.005063916993094608]}, 'Threaded_static': {"(['shape(100, 100)', 'shape(5, 5)'], {})": [250000, 0.00018066700431518257, 0.0001790839887689799, 0.00020249999943189323], "(['shape(500, 500)', 'shape(11, 11)'], {})": [30250000, 0.005990457982989028, 0.006161083991173655, 0.006471291999332607]}, 'Unthreaded': {"(['shape(100, 100)', 'shape(5, 5)'], {})": [250000, 0.00021341597312130034, 0.00020179100101813674, 0.00020345899974927306], "(['shape(500, 500)', 'shape(11, 11)'], {})": [30250000, 0.031798416981473565, 0.031586291996063665, 0.031558375019812956]}}
__pyx_t_3 = __Pyx_PyDict_NewPresized(6); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_250000); __Pyx_GIVEREF(__pyx_int_250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_009790166019229218); __Pyx_GIVEREF(__pyx_float_0_009790166019229218); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_009790166019229218)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_003298707975773141); __Pyx_GIVEREF(__pyx_float_0_003298707975773141); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_003298707975773141)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_003652999992482364); __Pyx_GIVEREF(__pyx_float_0_003652999992482364); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_003652999992482364)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_shape_5_5, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_30250000); __Pyx_GIVEREF(__pyx_int_30250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_30250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_007485667010769248); __Pyx_GIVEREF(__pyx_float_0_007485667010769248); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_007485667010769248)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_007351625012233853); __Pyx_GIVEREF(__pyx_float_0_007351625012233853); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_007351625012233853)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_007932541979243979); __Pyx_GIVEREF(__pyx_float_0_007932541979243979); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_007932541979243979)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_shape_11_11, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_OpenCL, __pyx_t_1) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_250000); __Pyx_GIVEREF(__pyx_int_250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00022562500089406967); __Pyx_GIVEREF(__pyx_float_0_00022562500089406967); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_00022562500089406967)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00018245799583382905); __Pyx_GIVEREF(__pyx_float_0_00018245799583382905); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_00018245799583382905)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00017095799557864666); __Pyx_GIVEREF(__pyx_float_0_00017095799557864666); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_00017095799557864666)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_shape_5_5, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_30250000); __Pyx_GIVEREF(__pyx_int_30250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_30250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_007719500019447878); __Pyx_GIVEREF(__pyx_float_0_007719500019447878); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_007719500019447878)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0060990000201854855); __Pyx_GIVEREF(__pyx_float_0_0060990000201854855); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_0060990000201854855)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_008289082994451746); __Pyx_GIVEREF(__pyx_float_0_008289082994451746); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_008289082994451746)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_shape_11_11, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Threaded, __pyx_t_1) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_250000); __Pyx_GIVEREF(__pyx_int_250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0001590839819982648); __Pyx_GIVEREF(__pyx_float_0_0001590839819982648); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_0001590839819982648)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0001595419889781624); __Pyx_GIVEREF(__pyx_float_0_0001595419889781624); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_0001595419889781624)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00020799998310394585); __Pyx_GIVEREF(__pyx_float_0_00020799998310394585); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_00020799998310394585)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_shape_5_5, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_30250000); __Pyx_GIVEREF(__pyx_int_30250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_30250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_005388707999372855); __Pyx_GIVEREF(__pyx_float_0_005388707999372855); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_005388707999372855)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_005681082984665409); __Pyx_GIVEREF(__pyx_float_0_005681082984665409); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_005681082984665409)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_005188624985748902); __Pyx_GIVEREF(__pyx_float_0_005188624985748902); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_005188624985748902)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_shape_11_11, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Threaded_dynamic, __pyx_t_1) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_250000); __Pyx_GIVEREF(__pyx_int_250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00018945799092762172); __Pyx_GIVEREF(__pyx_float_0_00018945799092762172); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_00018945799092762172)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0001723750028759241); __Pyx_GIVEREF(__pyx_float_0_0001723750028759241); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_0001723750028759241)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00017620899598114192); __Pyx_GIVEREF(__pyx_float_0_00017620899598114192); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_00017620899598114192)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_shape_5_5, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_30250000); __Pyx_GIVEREF(__pyx_int_30250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_30250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_005136792024131864); __Pyx_GIVEREF(__pyx_float_0_005136792024131864); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_005136792024131864)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_005925459001446143); __Pyx_GIVEREF(__pyx_float_0_005925459001446143); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_005925459001446143)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_005063916993094608); __Pyx_GIVEREF(__pyx_float_0_005063916993094608); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_005063916993094608)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_shape_11_11, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Threaded_guided, __pyx_t_1) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_250000); __Pyx_GIVEREF(__pyx_int_250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00018066700431518257); __Pyx_GIVEREF(__pyx_float_0_00018066700431518257); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_00018066700431518257)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0001790839887689799); __Pyx_GIVEREF(__pyx_float_0_0001790839887689799); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_0001790839887689799)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00020249999943189323); __Pyx_GIVEREF(__pyx_float_0_00020249999943189323); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_00020249999943189323)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_shape_5_5, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_30250000); __Pyx_GIVEREF(__pyx_int_30250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_30250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_005990457982989028); __Pyx_GIVEREF(__pyx_float_0_005990457982989028); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_005990457982989028)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_006161083991173655); __Pyx_GIVEREF(__pyx_float_0_006161083991173655); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_006161083991173655)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_006471291999332607); __Pyx_GIVEREF(__pyx_float_0_006471291999332607); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_006471291999332607)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_shape_11_11, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Threaded_static, __pyx_t_1) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_250000); __Pyx_GIVEREF(__pyx_int_250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00021341597312130034); __Pyx_GIVEREF(__pyx_float_0_00021341597312130034); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_00021341597312130034)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00020179100101813674); __Pyx_GIVEREF(__pyx_float_0_00020179100101813674); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_00020179100101813674)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_00020345899974927306); __Pyx_GIVEREF(__pyx_float_0_00020345899974927306); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_00020345899974927306)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_shape_5_5, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_30250000); __Pyx_GIVEREF(__pyx_int_30250000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_30250000)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_031798416981473565); __Pyx_GIVEREF(__pyx_float_0_031798416981473565); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_031798416981473565)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_031586291996063665); __Pyx_GIVEREF(__pyx_float_0_031586291996063665); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_031586291996063665)) __PYX_ERR(0, 28, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_031558375019812956); __Pyx_GIVEREF(__pyx_float_0_031558375019812956); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_031558375019812956)) __PYX_ERR(0, 28, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_shape_11_11, __pyx_t_2) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Unthreaded, __pyx_t_1) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; if (__Pyx_PyObject_SetAttrStr(__pyx_v_self, __pyx_n_s_default_benchmarks, __pyx_t_3) < 0) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0;
029:
+030: def run(self, image, kernel, run_type=None):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_3run(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_3run = {"run", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_3run, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_3run(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
PyObject *__pyx_v_self = 0;
PyObject *__pyx_v_image = 0;
PyObject *__pyx_v_kernel = 0;
PyObject *__pyx_v_run_type = 0;
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("run (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 30, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_kernel,&__pyx_n_s_run_type,0};
PyObject* values[4] = {0,0,0,0};
values[3] = __Pyx_Arg_NewRef_FASTCALL(((PyObject *)Py_None));
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 4: values[3] = __Pyx_Arg_FASTCALL(__pyx_args, 3);
CYTHON_FALLTHROUGH;
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 30, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 30, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("run", 0, 3, 4, 1); __PYX_ERR(0, 30, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (likely((values[2] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_kernel)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[2]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 30, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("run", 0, 3, 4, 2); __PYX_ERR(0, 30, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 3:
if (kw_args > 0) {
PyObject* value = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_run_type);
if (value) { values[3] = __Pyx_Arg_NewRef_FASTCALL(value); kw_args--; }
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 30, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "run") < 0)) __PYX_ERR(0, 30, __pyx_L3_error)
}
} else {
switch (__pyx_nargs) {
case 4: values[3] = __Pyx_Arg_FASTCALL(__pyx_args, 3);
CYTHON_FALLTHROUGH;
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
break;
default: goto __pyx_L5_argtuple_error;
}
}
__pyx_v_self = values[0];
__pyx_v_image = values[1];
__pyx_v_kernel = values[2];
__pyx_v_run_type = values[3];
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("run", 0, 3, 4, __pyx_nargs); __PYX_ERR(0, 30, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution.run", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_2run(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_kernel, __pyx_v_run_type);
/* function exit code */
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_2run(CYTHON_UNUSED PyObject *__pyx_self, PyObject *__pyx_v_self, PyObject *__pyx_v_image, PyObject *__pyx_v_kernel, PyObject *__pyx_v_run_type) {
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("run", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution.run", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__26 = PyTuple_Pack(4, __pyx_n_s_self, __pyx_n_s_image, __pyx_n_s_kernel, __pyx_n_s_run_type); if (unlikely(!__pyx_tuple__26)) __PYX_ERR(0, 30, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__26);
__Pyx_GIVEREF(__pyx_tuple__26);
__pyx_codeobj__27 = (PyObject*)__Pyx_PyCode_New(4, 0, 0, 4, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__26, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_convoluti, __pyx_n_s_run_2, 30, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__27)) __PYX_ERR(0, 30, __pyx_L1_error)
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_3run, 0, __pyx_n_s_Convolution_run, NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_d, ((PyObject *)__pyx_codeobj__27)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 30, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
__Pyx_CyFunction_SetDefaultsTuple(__pyx_t_10, __pyx_tuple__28);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_2, __pyx_t_10) < 0) __PYX_ERR(0, 30, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
__pyx_tuple__28 = PyTuple_Pack(1, Py_None); if (unlikely(!__pyx_tuple__28)) __PYX_ERR(0, 30, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__28);
__Pyx_GIVEREF(__pyx_tuple__28);
+031: return self._run(image, kernel, run_type=run_type)
__Pyx_XDECREF(__pyx_r); __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_v_self, __pyx_n_s_run); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 31, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyTuple_New(2); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 31, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_v_image); __Pyx_GIVEREF(__pyx_v_image); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_2, 0, __pyx_v_image)) __PYX_ERR(0, 31, __pyx_L1_error); __Pyx_INCREF(__pyx_v_kernel); __Pyx_GIVEREF(__pyx_v_kernel); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_2, 1, __pyx_v_kernel)) __PYX_ERR(0, 31, __pyx_L1_error); __pyx_t_3 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 31, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); if (PyDict_SetItem(__pyx_t_3, __pyx_n_s_run_type, __pyx_v_run_type) < 0) __PYX_ERR(0, 31, __pyx_L1_error) __pyx_t_4 = __Pyx_PyObject_Call(__pyx_t_1, __pyx_t_2, __pyx_t_3); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 31, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_r = __pyx_t_4; __pyx_t_4 = 0; goto __pyx_L0;
032:
+033: def benchmark(self, image, kernel):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_5benchmark(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_5benchmark = {"benchmark", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_5benchmark, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_5benchmark(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
PyObject *__pyx_v_self = 0;
PyObject *__pyx_v_image = 0;
PyObject *__pyx_v_kernel = 0;
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("benchmark (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 33, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_kernel,0};
PyObject* values[3] = {0,0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 33, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 33, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("benchmark", 1, 3, 3, 1); __PYX_ERR(0, 33, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (likely((values[2] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_kernel)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[2]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 33, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("benchmark", 1, 3, 3, 2); __PYX_ERR(0, 33, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "benchmark") < 0)) __PYX_ERR(0, 33, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 3)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
}
__pyx_v_self = values[0];
__pyx_v_image = values[1];
__pyx_v_kernel = values[2];
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("benchmark", 1, 3, 3, __pyx_nargs); __PYX_ERR(0, 33, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution.benchmark", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_4benchmark(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_kernel);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_4benchmark(CYTHON_UNUSED PyObject *__pyx_self, PyObject *__pyx_v_self, PyObject *__pyx_v_image, PyObject *__pyx_v_kernel) {
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("benchmark", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution.benchmark", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__29 = PyTuple_Pack(3, __pyx_n_s_self, __pyx_n_s_image, __pyx_n_s_kernel); if (unlikely(!__pyx_tuple__29)) __PYX_ERR(0, 33, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__29);
__Pyx_GIVEREF(__pyx_tuple__29);
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_5benchmark, 0, __pyx_n_s_Convolution_benchmark, NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_d, ((PyObject *)__pyx_codeobj__30)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 33, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
__Pyx_INCREF(__pyx_t_10);
PyList_Append(__pyx_t_7, __pyx_t_10);
__Pyx_GIVEREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_benchmark, __pyx_t_10) < 0) __PYX_ERR(0, 33, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
__pyx_codeobj__30 = (PyObject*)__Pyx_PyCode_New(3, 0, 0, 3, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__29, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_convoluti, __pyx_n_s_benchmark, 33, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__30)) __PYX_ERR(0, 33, __pyx_L1_error)
+034: return super().benchmark(image, kernel)
__Pyx_XDECREF(__pyx_r); __pyx_t_2 = __Pyx_CyFunction_GetClassObj(__pyx_self); if (!__pyx_t_2) { PyErr_SetString(PyExc_SystemError, "super(): empty __class__ cell"); __PYX_ERR(0, 34, __pyx_L1_error) } __Pyx_INCREF(__pyx_t_2); __pyx_t_3 = PyTuple_New(2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 34, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_2)) __PYX_ERR(0, 34, __pyx_L1_error); __Pyx_INCREF(__pyx_v_self); __Pyx_GIVEREF(__pyx_v_self); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 1, __pyx_v_self)) __PYX_ERR(0, 34, __pyx_L1_error); __pyx_t_2 = 0; __pyx_t_2 = __Pyx_PyObject_Call(__pyx_builtin_super, __pyx_t_3, NULL); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 34, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_benchmark); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 34, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = NULL; __pyx_t_4 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_3))) { __pyx_t_2 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_2)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_2); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_4 = 1; } } #endif { PyObject *__pyx_callargs[3] = {__pyx_t_2, __pyx_v_image, __pyx_v_kernel}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_4, 2+__pyx_t_4); __Pyx_XDECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 34, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_r = __pyx_t_1; __pyx_t_1 = 0; goto __pyx_L0;
035:
+036: def _run_unthreaded(self, float[:,:] image, float[:,:] kernel):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_7_run_unthreaded(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_7_run_unthreaded = {"_run_unthreaded", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_7_run_unthreaded, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_7_run_unthreaded(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_kernel = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_unthreaded (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 36, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_kernel,0};
PyObject* values[3] = {0,0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 36, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 36, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_unthreaded", 1, 3, 3, 1); __PYX_ERR(0, 36, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (likely((values[2] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_kernel)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[2]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 36, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_unthreaded", 1, 3, 3, 2); __PYX_ERR(0, 36, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_unthreaded") < 0)) __PYX_ERR(0, 36, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 3)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 36, __pyx_L3_error)
__pyx_v_kernel = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[2], PyBUF_WRITABLE); if (unlikely(!__pyx_v_kernel.memview)) __PYX_ERR(0, 36, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_unthreaded", 1, 3, 3, __pyx_nargs); __PYX_ERR(0, 36, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_unthreaded", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_6_run_unthreaded(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_kernel);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_6_run_unthreaded(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image, __Pyx_memviewslice __pyx_v_kernel) {
int __pyx_v_nRows;
int __pyx_v_nCols;
int __pyx_v_nRows_kernel;
int __pyx_v_nCols_kernel;
int __pyx_v_center_r;
int __pyx_v_center_c;
int __pyx_v_r;
int __pyx_v_c;
int __pyx_v_kr;
int __pyx_v_kc;
int __pyx_v_local_row;
int __pyx_v_local_col;
float __pyx_v_acc;
PyObject *__pyx_v_conv_out = NULL;
__Pyx_memviewslice __pyx_v__conv_out = { 0, 0, { 0 }, { 0 }, { 0 } };
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_unthreaded", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__Pyx_XDECREF(__pyx_t_5);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_unthreaded", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XDECREF(__pyx_v_conv_out);
__PYX_XCLEAR_MEMVIEW(&__pyx_v__conv_out, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__31 = PyTuple_Pack(18, __pyx_n_s_self, __pyx_n_s_image, __pyx_n_s_kernel, __pyx_n_s_nRows, __pyx_n_s_nCols, __pyx_n_s_nRows_kernel, __pyx_n_s_nCols_kernel, __pyx_n_s_center_r, __pyx_n_s_center_c, __pyx_n_s_r, __pyx_n_s_c, __pyx_n_s_kr, __pyx_n_s_kc, __pyx_n_s_local_row, __pyx_n_s_local_col, __pyx_n_s_acc, __pyx_n_s_conv_out, __pyx_n_s_conv_out_2); if (unlikely(!__pyx_tuple__31)) __PYX_ERR(0, 36, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__31);
__Pyx_GIVEREF(__pyx_tuple__31);
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_7_run_unthreaded, 0, __pyx_n_s_Convolution__run_unthreaded, NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_d, ((PyObject *)__pyx_codeobj__32)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 36, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_unthreaded, __pyx_t_10) < 0) __PYX_ERR(0, 36, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
__pyx_codeobj__32 = (PyObject*)__Pyx_PyCode_New(3, 0, 0, 18, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__31, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_convoluti, __pyx_n_s_run_unthreaded, 36, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__32)) __PYX_ERR(0, 36, __pyx_L1_error)
037:
+038: cdef int nRows = image.shape[0]
__pyx_v_nRows = (__pyx_v_image.shape[0]);
+039: cdef int nCols = image.shape[1]
__pyx_v_nCols = (__pyx_v_image.shape[1]);
040:
+041: cdef int nRows_kernel = kernel.shape[0]
__pyx_v_nRows_kernel = (__pyx_v_kernel.shape[0]);
+042: cdef int nCols_kernel = kernel.shape[1]
__pyx_v_nCols_kernel = (__pyx_v_kernel.shape[1]);
043:
+044: cdef int center_r = (nRows_kernel-1) // 2
__pyx_v_center_r = ((__pyx_v_nRows_kernel - 1) / 2);
+045: cdef int center_c = (nCols_kernel-1) // 2
__pyx_v_center_c = ((__pyx_v_nCols_kernel - 1) / 2);
046:
047: cdef int r,c
048: cdef int kr,kc
049:
050: cdef int local_row, local_col
051:
+052: cdef float acc = 0
__pyx_v_acc = 0.0;
053:
054:
+055: conv_out = np.zeros((nRows, nCols), dtype=np.float32)
__Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_zeros); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyInt_From_int(__pyx_v_nRows); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_3 = __Pyx_PyInt_From_int(__pyx_v_nCols); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = PyTuple_New(2); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 0, __pyx_t_1)) __PYX_ERR(0, 55, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_3); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 1, __pyx_t_3)) __PYX_ERR(0, 55, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_3 = 0; __pyx_t_3 = PyTuple_New(1); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_4); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_4)) __PYX_ERR(0, 55, __pyx_L1_error); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_5 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_float32); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; if (PyDict_SetItem(__pyx_t_4, __pyx_n_s_dtype, __pyx_t_5) < 0) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_DECREF(__pyx_t_5); __pyx_t_5 = 0; __pyx_t_5 = __Pyx_PyObject_Call(__pyx_t_2, __pyx_t_3, __pyx_t_4); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 55, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_v_conv_out = __pyx_t_5; __pyx_t_5 = 0;
+056: cdef float[:,:] _conv_out = conv_out
__pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(__pyx_v_conv_out, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 56, __pyx_L1_error) __pyx_v__conv_out = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
057:
+058: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+059: for r in range(nRows):
__pyx_t_7 = __pyx_v_nRows;
__pyx_t_8 = __pyx_t_7;
for (__pyx_t_9 = 0; __pyx_t_9 < __pyx_t_8; __pyx_t_9+=1) {
__pyx_v_r = __pyx_t_9;
+060: for c in range(nCols):
__pyx_t_10 = __pyx_v_nCols;
__pyx_t_11 = __pyx_t_10;
for (__pyx_t_12 = 0; __pyx_t_12 < __pyx_t_11; __pyx_t_12+=1) {
__pyx_v_c = __pyx_t_12;
+061: acc = 0
__pyx_v_acc = 0.0;
+062: for kr in range(nRows_kernel):
__pyx_t_13 = __pyx_v_nRows_kernel;
__pyx_t_14 = __pyx_t_13;
for (__pyx_t_15 = 0; __pyx_t_15 < __pyx_t_14; __pyx_t_15+=1) {
__pyx_v_kr = __pyx_t_15;
+063: for kc in range(nCols_kernel):
__pyx_t_16 = __pyx_v_nCols_kernel;
__pyx_t_17 = __pyx_t_16;
for (__pyx_t_18 = 0; __pyx_t_18 < __pyx_t_17; __pyx_t_18+=1) {
__pyx_v_kc = __pyx_t_18;
+064: local_row = min(max(r+(kr-center_r),0),nRows-1)
__pyx_t_19 = (__pyx_v_nRows - 1);
__pyx_t_20 = 0;
__pyx_t_21 = (__pyx_v_r + (__pyx_v_kr - __pyx_v_center_r));
__pyx_t_23 = (__pyx_t_20 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_20;
} else {
__pyx_t_22 = __pyx_t_21;
}
__pyx_t_20 = __pyx_t_22;
__pyx_t_23 = (__pyx_t_19 < __pyx_t_20);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_19;
} else {
__pyx_t_22 = __pyx_t_20;
}
__pyx_v_local_row = __pyx_t_22;
+065: local_col = min(max(c+(kc-center_c),0),nCols-1)
__pyx_t_22 = (__pyx_v_nCols - 1);
__pyx_t_19 = 0;
__pyx_t_21 = (__pyx_v_c + (__pyx_v_kc - __pyx_v_center_c));
__pyx_t_23 = (__pyx_t_19 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_19;
} else {
__pyx_t_20 = __pyx_t_21;
}
__pyx_t_19 = __pyx_t_20;
__pyx_t_23 = (__pyx_t_22 < __pyx_t_19);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_22;
} else {
__pyx_t_20 = __pyx_t_19;
}
__pyx_v_local_col = __pyx_t_20;
+066: acc = acc + kernel[kr,kc] * image[local_row, local_col]
__pyx_t_24 = __pyx_v_kr;
__pyx_t_25 = __pyx_v_kc;
__pyx_t_26 = __pyx_v_local_row;
__pyx_t_27 = __pyx_v_local_col;
__pyx_v_acc = (__pyx_v_acc + ((*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_kernel.data + __pyx_t_24 * __pyx_v_kernel.strides[0]) ) + __pyx_t_25 * __pyx_v_kernel.strides[1]) ))) * (*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_26 * __pyx_v_image.strides[0]) ) + __pyx_t_27 * __pyx_v_image.strides[1]) )))));
}
}
+067: _conv_out[r,c] = acc
__pyx_t_27 = __pyx_v_r;
__pyx_t_26 = __pyx_v_c;
*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v__conv_out.data + __pyx_t_27 * __pyx_v__conv_out.strides[0]) ) + __pyx_t_26 * __pyx_v__conv_out.strides[1]) )) = __pyx_v_acc;
}
}
}
068:
+069: return conv_out
__Pyx_XDECREF(__pyx_r); __Pyx_INCREF(__pyx_v_conv_out); __pyx_r = __pyx_v_conv_out; goto __pyx_L0;
070:
+071: def _run_threaded(self, float[:,:] image, float[:,:] kernel):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_9_run_threaded(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_9_run_threaded = {"_run_threaded", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_9_run_threaded, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_9_run_threaded(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_kernel = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 71, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_kernel,0};
PyObject* values[3] = {0,0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 71, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 71, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded", 1, 3, 3, 1); __PYX_ERR(0, 71, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (likely((values[2] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_kernel)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[2]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 71, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded", 1, 3, 3, 2); __PYX_ERR(0, 71, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_threaded") < 0)) __PYX_ERR(0, 71, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 3)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 71, __pyx_L3_error)
__pyx_v_kernel = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[2], PyBUF_WRITABLE); if (unlikely(!__pyx_v_kernel.memview)) __PYX_ERR(0, 71, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_threaded", 1, 3, 3, __pyx_nargs); __PYX_ERR(0, 71, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_threaded", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_8_run_threaded(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_kernel);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_8_run_threaded(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image, __Pyx_memviewslice __pyx_v_kernel) {
int __pyx_v_nRows;
int __pyx_v_nCols;
int __pyx_v_nRows_kernel;
int __pyx_v_nCols_kernel;
int __pyx_v_center_r;
int __pyx_v_center_c;
int __pyx_v_r;
int __pyx_v_c;
int __pyx_v_kr;
int __pyx_v_kc;
int __pyx_v_local_row;
int __pyx_v_local_col;
float __pyx_v_acc;
PyObject *__pyx_v_conv_out = NULL;
__Pyx_memviewslice __pyx_v__conv_out = { 0, 0, { 0 }, { 0 }, { 0 } };
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__Pyx_XDECREF(__pyx_t_5);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_threaded", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XDECREF(__pyx_v_conv_out);
__PYX_XCLEAR_MEMVIEW(&__pyx_v__conv_out, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_9_run_threaded, 0, __pyx_n_s_Convolution__run_threaded, NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_d, ((PyObject *)__pyx_codeobj__33)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 71, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_threaded, __pyx_t_10) < 0) __PYX_ERR(0, 71, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
/* … */
__pyx_codeobj__33 = (PyObject*)__Pyx_PyCode_New(3, 0, 0, 18, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__31, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_convoluti, __pyx_n_s_run_threaded, 71, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__33)) __PYX_ERR(0, 71, __pyx_L1_error)
072:
+073: cdef int nRows = image.shape[0]
__pyx_v_nRows = (__pyx_v_image.shape[0]);
+074: cdef int nCols = image.shape[1]
__pyx_v_nCols = (__pyx_v_image.shape[1]);
075:
+076: cdef int nRows_kernel = kernel.shape[0]
__pyx_v_nRows_kernel = (__pyx_v_kernel.shape[0]);
+077: cdef int nCols_kernel = kernel.shape[1]
__pyx_v_nCols_kernel = (__pyx_v_kernel.shape[1]);
078:
+079: cdef int center_r = (nRows_kernel-1) // 2
__pyx_v_center_r = ((__pyx_v_nRows_kernel - 1) / 2);
+080: cdef int center_c = (nCols_kernel-1) // 2
__pyx_v_center_c = ((__pyx_v_nCols_kernel - 1) / 2);
081:
082: cdef int r,c
083: cdef int kr,kc
084:
085: cdef int local_row, local_col
086:
+087: cdef float acc = 0
__pyx_v_acc = 0.0;
088:
+089: conv_out = np.zeros((nRows, nCols), dtype=np.float32)
__Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_zeros); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyInt_From_int(__pyx_v_nRows); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_3 = __Pyx_PyInt_From_int(__pyx_v_nCols); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = PyTuple_New(2); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 0, __pyx_t_1)) __PYX_ERR(0, 89, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_3); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 1, __pyx_t_3)) __PYX_ERR(0, 89, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_3 = 0; __pyx_t_3 = PyTuple_New(1); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_4); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_4)) __PYX_ERR(0, 89, __pyx_L1_error); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_5 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_float32); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; if (PyDict_SetItem(__pyx_t_4, __pyx_n_s_dtype, __pyx_t_5) < 0) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_DECREF(__pyx_t_5); __pyx_t_5 = 0; __pyx_t_5 = __Pyx_PyObject_Call(__pyx_t_2, __pyx_t_3, __pyx_t_4); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 89, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_v_conv_out = __pyx_t_5; __pyx_t_5 = 0;
+090: cdef float[:,:] _conv_out = conv_out
__pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(__pyx_v_conv_out, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 90, __pyx_L1_error) __pyx_v__conv_out = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
091:
+092: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+093: for r in prange(nRows):
__pyx_t_7 = __pyx_v_nRows;
{
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) (x)
#define unlikely(x) (x)
#endif
__pyx_t_9 = (__pyx_t_7 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_9 > 0)
{
#ifdef _OPENMP
#pragma omp parallel
#endif /* _OPENMP */
{
#ifdef _OPENMP
#pragma omp for lastprivate(__pyx_v_acc) lastprivate(__pyx_v_c) lastprivate(__pyx_v_kc) lastprivate(__pyx_v_kr) lastprivate(__pyx_v_local_col) lastprivate(__pyx_v_local_row) firstprivate(__pyx_v_r) lastprivate(__pyx_v_r)
#endif /* _OPENMP */
for (__pyx_t_8 = 0; __pyx_t_8 < __pyx_t_9; __pyx_t_8++){
{
__pyx_v_r = (int)(0 + 1 * __pyx_t_8);
/* Initialize private variables to invalid values */
__pyx_v_acc = ((float)__PYX_NAN());
__pyx_v_c = ((int)0xbad0bad0);
__pyx_v_kc = ((int)0xbad0bad0);
__pyx_v_kr = ((int)0xbad0bad0);
__pyx_v_local_col = ((int)0xbad0bad0);
__pyx_v_local_row = ((int)0xbad0bad0);
+094: for c in prange(nCols):
__pyx_t_10 = __pyx_v_nCols;
{
__pyx_t_12 = (__pyx_t_10 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_12 > 0)
{
#if 0
#pragma omp parallel
#endif /* _OPENMP */
{
#if 0
#pragma omp for lastprivate(__pyx_v_acc) firstprivate(__pyx_v_c) lastprivate(__pyx_v_c) lastprivate(__pyx_v_kc) lastprivate(__pyx_v_kr) lastprivate(__pyx_v_local_col) lastprivate(__pyx_v_local_row)
#endif /* _OPENMP */
for (__pyx_t_11 = 0; __pyx_t_11 < __pyx_t_12; __pyx_t_11++){
{
__pyx_v_c = (int)(0 + 1 * __pyx_t_11);
/* Initialize private variables to invalid values */
__pyx_v_acc = ((float)__PYX_NAN());
__pyx_v_kc = ((int)0xbad0bad0);
__pyx_v_kr = ((int)0xbad0bad0);
__pyx_v_local_col = ((int)0xbad0bad0);
__pyx_v_local_row = ((int)0xbad0bad0);
+095: acc = 0
__pyx_v_acc = 0.0;
+096: for kr in range(nRows_kernel):
__pyx_t_13 = __pyx_v_nRows_kernel;
__pyx_t_14 = __pyx_t_13;
for (__pyx_t_15 = 0; __pyx_t_15 < __pyx_t_14; __pyx_t_15+=1) {
__pyx_v_kr = __pyx_t_15;
+097: for kc in range(nCols_kernel):
__pyx_t_16 = __pyx_v_nCols_kernel;
__pyx_t_17 = __pyx_t_16;
for (__pyx_t_18 = 0; __pyx_t_18 < __pyx_t_17; __pyx_t_18+=1) {
__pyx_v_kc = __pyx_t_18;
+098: local_row = min(max(r+(kr-center_r),0),nRows-1)
__pyx_t_19 = (__pyx_v_nRows - 1);
__pyx_t_20 = 0;
__pyx_t_21 = (__pyx_v_r + (__pyx_v_kr - __pyx_v_center_r));
__pyx_t_23 = (__pyx_t_20 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_20;
} else {
__pyx_t_22 = __pyx_t_21;
}
__pyx_t_20 = __pyx_t_22;
__pyx_t_23 = (__pyx_t_19 < __pyx_t_20);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_19;
} else {
__pyx_t_22 = __pyx_t_20;
}
__pyx_v_local_row = __pyx_t_22;
+099: local_col = min(max(c+(kc-center_c),0),nCols-1)
__pyx_t_22 = (__pyx_v_nCols - 1);
__pyx_t_19 = 0;
__pyx_t_21 = (__pyx_v_c + (__pyx_v_kc - __pyx_v_center_c));
__pyx_t_23 = (__pyx_t_19 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_19;
} else {
__pyx_t_20 = __pyx_t_21;
}
__pyx_t_19 = __pyx_t_20;
__pyx_t_23 = (__pyx_t_22 < __pyx_t_19);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_22;
} else {
__pyx_t_20 = __pyx_t_19;
}
__pyx_v_local_col = __pyx_t_20;
+100: acc = acc + kernel[kr,kc] * image[local_row, local_col]
__pyx_t_24 = __pyx_v_kr;
__pyx_t_25 = __pyx_v_kc;
__pyx_t_26 = __pyx_v_local_row;
__pyx_t_27 = __pyx_v_local_col;
__pyx_v_acc = (__pyx_v_acc + ((*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_kernel.data + __pyx_t_24 * __pyx_v_kernel.strides[0]) ) + __pyx_t_25 * __pyx_v_kernel.strides[1]) ))) * (*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_26 * __pyx_v_image.strides[0]) ) + __pyx_t_27 * __pyx_v_image.strides[1]) )))));
}
}
+101: _conv_out[r,c] = acc
__pyx_t_27 = __pyx_v_r;
__pyx_t_26 = __pyx_v_c;
*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v__conv_out.data + __pyx_t_27 * __pyx_v__conv_out.strides[0]) ) + __pyx_t_26 * __pyx_v__conv_out.strides[1]) )) = __pyx_v_acc;
}
}
}
}
}
}
}
}
}
}
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) __builtin_expect(!!(x), 1)
#define unlikely(x) __builtin_expect(!!(x), 0)
#endif
}
102:
+103: return conv_out
__Pyx_XDECREF(__pyx_r); __Pyx_INCREF(__pyx_v_conv_out); __pyx_r = __pyx_v_conv_out; goto __pyx_L0;
104:
+105: def _run_threaded_static(self, float[:,:] image, float[:,:] kernel):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_11_run_threaded_static(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_11_run_threaded_static = {"_run_threaded_static", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_11_run_threaded_static, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_11_run_threaded_static(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_kernel = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_static (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 105, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_kernel,0};
PyObject* values[3] = {0,0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 105, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 105, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded_static", 1, 3, 3, 1); __PYX_ERR(0, 105, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (likely((values[2] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_kernel)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[2]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 105, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded_static", 1, 3, 3, 2); __PYX_ERR(0, 105, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_threaded_static") < 0)) __PYX_ERR(0, 105, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 3)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 105, __pyx_L3_error)
__pyx_v_kernel = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[2], PyBUF_WRITABLE); if (unlikely(!__pyx_v_kernel.memview)) __PYX_ERR(0, 105, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_threaded_static", 1, 3, 3, __pyx_nargs); __PYX_ERR(0, 105, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_threaded_static", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_10_run_threaded_static(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_kernel);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_10_run_threaded_static(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image, __Pyx_memviewslice __pyx_v_kernel) {
int __pyx_v_nRows;
int __pyx_v_nCols;
int __pyx_v_nRows_kernel;
int __pyx_v_nCols_kernel;
int __pyx_v_center_r;
int __pyx_v_center_c;
int __pyx_v_r;
int __pyx_v_c;
int __pyx_v_kr;
int __pyx_v_kc;
int __pyx_v_local_row;
int __pyx_v_local_col;
float __pyx_v_acc;
PyObject *__pyx_v_conv_out = NULL;
__Pyx_memviewslice __pyx_v__conv_out = { 0, 0, { 0 }, { 0 }, { 0 } };
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_static", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__Pyx_XDECREF(__pyx_t_5);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_threaded_static", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XDECREF(__pyx_v_conv_out);
__PYX_XCLEAR_MEMVIEW(&__pyx_v__conv_out, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_11_run_threaded_static, 0, __pyx_n_s_Convolution__run_threaded_static, NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_d, ((PyObject *)__pyx_codeobj__34)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 105, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_threaded_static, __pyx_t_10) < 0) __PYX_ERR(0, 105, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
/* … */
__pyx_codeobj__34 = (PyObject*)__Pyx_PyCode_New(3, 0, 0, 18, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__31, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_convoluti, __pyx_n_s_run_threaded_static, 105, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__34)) __PYX_ERR(0, 105, __pyx_L1_error)
106:
+107: cdef int nRows = image.shape[0]
__pyx_v_nRows = (__pyx_v_image.shape[0]);
+108: cdef int nCols = image.shape[1]
__pyx_v_nCols = (__pyx_v_image.shape[1]);
109:
+110: cdef int nRows_kernel = kernel.shape[0]
__pyx_v_nRows_kernel = (__pyx_v_kernel.shape[0]);
+111: cdef int nCols_kernel = kernel.shape[1]
__pyx_v_nCols_kernel = (__pyx_v_kernel.shape[1]);
112:
+113: cdef int center_r = (nRows_kernel-1) // 2
__pyx_v_center_r = ((__pyx_v_nRows_kernel - 1) / 2);
+114: cdef int center_c = (nCols_kernel-1) // 2
__pyx_v_center_c = ((__pyx_v_nCols_kernel - 1) / 2);
115:
116: cdef int r,c
117: cdef int kr,kc
118:
119: cdef int local_row, local_col
120:
+121: cdef float acc = 0
__pyx_v_acc = 0.0;
122:
+123: conv_out = np.zeros((nRows, nCols), dtype=np.float32)
__Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_zeros); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyInt_From_int(__pyx_v_nRows); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_3 = __Pyx_PyInt_From_int(__pyx_v_nCols); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = PyTuple_New(2); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 0, __pyx_t_1)) __PYX_ERR(0, 123, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_3); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 1, __pyx_t_3)) __PYX_ERR(0, 123, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_3 = 0; __pyx_t_3 = PyTuple_New(1); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_4); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_4)) __PYX_ERR(0, 123, __pyx_L1_error); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_5 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_float32); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; if (PyDict_SetItem(__pyx_t_4, __pyx_n_s_dtype, __pyx_t_5) < 0) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_DECREF(__pyx_t_5); __pyx_t_5 = 0; __pyx_t_5 = __Pyx_PyObject_Call(__pyx_t_2, __pyx_t_3, __pyx_t_4); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 123, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_v_conv_out = __pyx_t_5; __pyx_t_5 = 0;
+124: cdef float[:,:] _conv_out = conv_out
__pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(__pyx_v_conv_out, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 124, __pyx_L1_error) __pyx_v__conv_out = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
125:
+126: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+127: for r in prange(nRows, schedule="static"):
__pyx_t_7 = __pyx_v_nRows;
{
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) (x)
#define unlikely(x) (x)
#endif
__pyx_t_9 = (__pyx_t_7 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_9 > 0)
{
#ifdef _OPENMP
#pragma omp parallel
#endif /* _OPENMP */
{
#ifdef _OPENMP
#pragma omp for lastprivate(__pyx_v_acc) lastprivate(__pyx_v_c) lastprivate(__pyx_v_kc) lastprivate(__pyx_v_kr) lastprivate(__pyx_v_local_col) lastprivate(__pyx_v_local_row) firstprivate(__pyx_v_r) lastprivate(__pyx_v_r) schedule(static)
#endif /* _OPENMP */
for (__pyx_t_8 = 0; __pyx_t_8 < __pyx_t_9; __pyx_t_8++){
{
__pyx_v_r = (int)(0 + 1 * __pyx_t_8);
/* Initialize private variables to invalid values */
__pyx_v_acc = ((float)__PYX_NAN());
__pyx_v_c = ((int)0xbad0bad0);
__pyx_v_kc = ((int)0xbad0bad0);
__pyx_v_kr = ((int)0xbad0bad0);
__pyx_v_local_col = ((int)0xbad0bad0);
__pyx_v_local_row = ((int)0xbad0bad0);
+128: for c in prange(nCols, schedule="static"):
__pyx_t_10 = __pyx_v_nCols;
{
__pyx_t_12 = (__pyx_t_10 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_12 > 0)
{
#if 0
#pragma omp parallel
#endif /* _OPENMP */
{
#if 0
#pragma omp for lastprivate(__pyx_v_acc) firstprivate(__pyx_v_c) lastprivate(__pyx_v_c) lastprivate(__pyx_v_kc) lastprivate(__pyx_v_kr) lastprivate(__pyx_v_local_col) lastprivate(__pyx_v_local_row) schedule(static)
#endif /* _OPENMP */
for (__pyx_t_11 = 0; __pyx_t_11 < __pyx_t_12; __pyx_t_11++){
{
__pyx_v_c = (int)(0 + 1 * __pyx_t_11);
/* Initialize private variables to invalid values */
__pyx_v_acc = ((float)__PYX_NAN());
__pyx_v_kc = ((int)0xbad0bad0);
__pyx_v_kr = ((int)0xbad0bad0);
__pyx_v_local_col = ((int)0xbad0bad0);
__pyx_v_local_row = ((int)0xbad0bad0);
+129: acc = 0
__pyx_v_acc = 0.0;
+130: for kr in range(nRows_kernel):
__pyx_t_13 = __pyx_v_nRows_kernel;
__pyx_t_14 = __pyx_t_13;
for (__pyx_t_15 = 0; __pyx_t_15 < __pyx_t_14; __pyx_t_15+=1) {
__pyx_v_kr = __pyx_t_15;
+131: for kc in range(nCols_kernel):
__pyx_t_16 = __pyx_v_nCols_kernel;
__pyx_t_17 = __pyx_t_16;
for (__pyx_t_18 = 0; __pyx_t_18 < __pyx_t_17; __pyx_t_18+=1) {
__pyx_v_kc = __pyx_t_18;
+132: local_row = min(max(r+(kr-center_r),0),nRows-1)
__pyx_t_19 = (__pyx_v_nRows - 1);
__pyx_t_20 = 0;
__pyx_t_21 = (__pyx_v_r + (__pyx_v_kr - __pyx_v_center_r));
__pyx_t_23 = (__pyx_t_20 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_20;
} else {
__pyx_t_22 = __pyx_t_21;
}
__pyx_t_20 = __pyx_t_22;
__pyx_t_23 = (__pyx_t_19 < __pyx_t_20);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_19;
} else {
__pyx_t_22 = __pyx_t_20;
}
__pyx_v_local_row = __pyx_t_22;
+133: local_col = min(max(c+(kc-center_c),0),nCols-1)
__pyx_t_22 = (__pyx_v_nCols - 1);
__pyx_t_19 = 0;
__pyx_t_21 = (__pyx_v_c + (__pyx_v_kc - __pyx_v_center_c));
__pyx_t_23 = (__pyx_t_19 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_19;
} else {
__pyx_t_20 = __pyx_t_21;
}
__pyx_t_19 = __pyx_t_20;
__pyx_t_23 = (__pyx_t_22 < __pyx_t_19);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_22;
} else {
__pyx_t_20 = __pyx_t_19;
}
__pyx_v_local_col = __pyx_t_20;
+134: acc = acc + kernel[kr,kc] * image[local_row, local_col]
__pyx_t_24 = __pyx_v_kr;
__pyx_t_25 = __pyx_v_kc;
__pyx_t_26 = __pyx_v_local_row;
__pyx_t_27 = __pyx_v_local_col;
__pyx_v_acc = (__pyx_v_acc + ((*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_kernel.data + __pyx_t_24 * __pyx_v_kernel.strides[0]) ) + __pyx_t_25 * __pyx_v_kernel.strides[1]) ))) * (*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_26 * __pyx_v_image.strides[0]) ) + __pyx_t_27 * __pyx_v_image.strides[1]) )))));
}
}
+135: _conv_out[r,c] = acc
__pyx_t_27 = __pyx_v_r;
__pyx_t_26 = __pyx_v_c;
*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v__conv_out.data + __pyx_t_27 * __pyx_v__conv_out.strides[0]) ) + __pyx_t_26 * __pyx_v__conv_out.strides[1]) )) = __pyx_v_acc;
}
}
}
}
}
}
}
}
}
}
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) __builtin_expect(!!(x), 1)
#define unlikely(x) __builtin_expect(!!(x), 0)
#endif
}
136:
+137: return conv_out
__Pyx_XDECREF(__pyx_r); __Pyx_INCREF(__pyx_v_conv_out); __pyx_r = __pyx_v_conv_out; goto __pyx_L0;
138:
+139: def _run_threaded_dynamic(self, float[:,:] image, float[:,:] kernel):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_13_run_threaded_dynamic(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_13_run_threaded_dynamic = {"_run_threaded_dynamic", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_13_run_threaded_dynamic, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_13_run_threaded_dynamic(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_kernel = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_dynamic (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 139, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_kernel,0};
PyObject* values[3] = {0,0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 139, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 139, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded_dynamic", 1, 3, 3, 1); __PYX_ERR(0, 139, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (likely((values[2] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_kernel)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[2]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 139, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded_dynamic", 1, 3, 3, 2); __PYX_ERR(0, 139, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_threaded_dynamic") < 0)) __PYX_ERR(0, 139, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 3)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 139, __pyx_L3_error)
__pyx_v_kernel = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[2], PyBUF_WRITABLE); if (unlikely(!__pyx_v_kernel.memview)) __PYX_ERR(0, 139, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_threaded_dynamic", 1, 3, 3, __pyx_nargs); __PYX_ERR(0, 139, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_threaded_dynamic", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_12_run_threaded_dynamic(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_kernel);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_12_run_threaded_dynamic(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image, __Pyx_memviewslice __pyx_v_kernel) {
int __pyx_v_nRows;
int __pyx_v_nCols;
int __pyx_v_nRows_kernel;
int __pyx_v_nCols_kernel;
int __pyx_v_center_r;
int __pyx_v_center_c;
int __pyx_v_r;
int __pyx_v_c;
int __pyx_v_kr;
int __pyx_v_kc;
int __pyx_v_local_row;
int __pyx_v_local_col;
float __pyx_v_acc;
PyObject *__pyx_v_conv_out = NULL;
__Pyx_memviewslice __pyx_v__conv_out = { 0, 0, { 0 }, { 0 }, { 0 } };
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_dynamic", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__Pyx_XDECREF(__pyx_t_5);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_threaded_dynamic", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XDECREF(__pyx_v_conv_out);
__PYX_XCLEAR_MEMVIEW(&__pyx_v__conv_out, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_13_run_threaded_dynamic, 0, __pyx_n_s_Convolution__run_threaded_dynami, NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_d, ((PyObject *)__pyx_codeobj__35)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 139, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_threaded_dynamic, __pyx_t_10) < 0) __PYX_ERR(0, 139, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
/* … */
__pyx_codeobj__35 = (PyObject*)__Pyx_PyCode_New(3, 0, 0, 18, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__31, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_convoluti, __pyx_n_s_run_threaded_dynamic, 139, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__35)) __PYX_ERR(0, 139, __pyx_L1_error)
140:
+141: cdef int nRows = image.shape[0]
__pyx_v_nRows = (__pyx_v_image.shape[0]);
+142: cdef int nCols = image.shape[1]
__pyx_v_nCols = (__pyx_v_image.shape[1]);
143:
+144: cdef int nRows_kernel = kernel.shape[0]
__pyx_v_nRows_kernel = (__pyx_v_kernel.shape[0]);
+145: cdef int nCols_kernel = kernel.shape[1]
__pyx_v_nCols_kernel = (__pyx_v_kernel.shape[1]);
146:
+147: cdef int center_r = (nRows_kernel-1) // 2
__pyx_v_center_r = ((__pyx_v_nRows_kernel - 1) / 2);
+148: cdef int center_c = (nCols_kernel-1) // 2
__pyx_v_center_c = ((__pyx_v_nCols_kernel - 1) / 2);
149:
150: cdef int r,c
151: cdef int kr,kc
152:
153: cdef int local_row, local_col
154:
+155: cdef float acc = 0
__pyx_v_acc = 0.0;
156:
+157: conv_out = np.zeros((nRows, nCols), dtype=np.float32)
__Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_zeros); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyInt_From_int(__pyx_v_nRows); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_3 = __Pyx_PyInt_From_int(__pyx_v_nCols); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = PyTuple_New(2); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 0, __pyx_t_1)) __PYX_ERR(0, 157, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_3); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 1, __pyx_t_3)) __PYX_ERR(0, 157, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_3 = 0; __pyx_t_3 = PyTuple_New(1); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_4); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_4)) __PYX_ERR(0, 157, __pyx_L1_error); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_5 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_float32); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; if (PyDict_SetItem(__pyx_t_4, __pyx_n_s_dtype, __pyx_t_5) < 0) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_DECREF(__pyx_t_5); __pyx_t_5 = 0; __pyx_t_5 = __Pyx_PyObject_Call(__pyx_t_2, __pyx_t_3, __pyx_t_4); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_v_conv_out = __pyx_t_5; __pyx_t_5 = 0;
+158: cdef float[:,:] _conv_out = conv_out
__pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(__pyx_v_conv_out, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 158, __pyx_L1_error) __pyx_v__conv_out = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
159:
+160: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+161: for r in prange(nRows, schedule="dynamic"):
__pyx_t_7 = __pyx_v_nRows;
{
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) (x)
#define unlikely(x) (x)
#endif
__pyx_t_9 = (__pyx_t_7 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_9 > 0)
{
#ifdef _OPENMP
#pragma omp parallel
#endif /* _OPENMP */
{
#ifdef _OPENMP
#pragma omp for lastprivate(__pyx_v_acc) lastprivate(__pyx_v_c) lastprivate(__pyx_v_kc) lastprivate(__pyx_v_kr) lastprivate(__pyx_v_local_col) lastprivate(__pyx_v_local_row) firstprivate(__pyx_v_r) lastprivate(__pyx_v_r) schedule(dynamic)
#endif /* _OPENMP */
for (__pyx_t_8 = 0; __pyx_t_8 < __pyx_t_9; __pyx_t_8++){
{
__pyx_v_r = (int)(0 + 1 * __pyx_t_8);
/* Initialize private variables to invalid values */
__pyx_v_acc = ((float)__PYX_NAN());
__pyx_v_c = ((int)0xbad0bad0);
__pyx_v_kc = ((int)0xbad0bad0);
__pyx_v_kr = ((int)0xbad0bad0);
__pyx_v_local_col = ((int)0xbad0bad0);
__pyx_v_local_row = ((int)0xbad0bad0);
+162: for c in prange(nCols, schedule="dynamic"):
__pyx_t_10 = __pyx_v_nCols;
{
__pyx_t_12 = (__pyx_t_10 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_12 > 0)
{
#if 0
#pragma omp parallel
#endif /* _OPENMP */
{
#if 0
#pragma omp for lastprivate(__pyx_v_acc) firstprivate(__pyx_v_c) lastprivate(__pyx_v_c) lastprivate(__pyx_v_kc) lastprivate(__pyx_v_kr) lastprivate(__pyx_v_local_col) lastprivate(__pyx_v_local_row) schedule(dynamic)
#endif /* _OPENMP */
for (__pyx_t_11 = 0; __pyx_t_11 < __pyx_t_12; __pyx_t_11++){
{
__pyx_v_c = (int)(0 + 1 * __pyx_t_11);
/* Initialize private variables to invalid values */
__pyx_v_acc = ((float)__PYX_NAN());
__pyx_v_kc = ((int)0xbad0bad0);
__pyx_v_kr = ((int)0xbad0bad0);
__pyx_v_local_col = ((int)0xbad0bad0);
__pyx_v_local_row = ((int)0xbad0bad0);
+163: acc = 0
__pyx_v_acc = 0.0;
+164: for kr in range(nRows_kernel):
__pyx_t_13 = __pyx_v_nRows_kernel;
__pyx_t_14 = __pyx_t_13;
for (__pyx_t_15 = 0; __pyx_t_15 < __pyx_t_14; __pyx_t_15+=1) {
__pyx_v_kr = __pyx_t_15;
+165: for kc in range(nCols_kernel):
__pyx_t_16 = __pyx_v_nCols_kernel;
__pyx_t_17 = __pyx_t_16;
for (__pyx_t_18 = 0; __pyx_t_18 < __pyx_t_17; __pyx_t_18+=1) {
__pyx_v_kc = __pyx_t_18;
+166: local_row = min(max(r+(kr-center_r),0),nRows-1)
__pyx_t_19 = (__pyx_v_nRows - 1);
__pyx_t_20 = 0;
__pyx_t_21 = (__pyx_v_r + (__pyx_v_kr - __pyx_v_center_r));
__pyx_t_23 = (__pyx_t_20 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_20;
} else {
__pyx_t_22 = __pyx_t_21;
}
__pyx_t_20 = __pyx_t_22;
__pyx_t_23 = (__pyx_t_19 < __pyx_t_20);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_19;
} else {
__pyx_t_22 = __pyx_t_20;
}
__pyx_v_local_row = __pyx_t_22;
+167: local_col = min(max(c+(kc-center_c),0),nCols-1)
__pyx_t_22 = (__pyx_v_nCols - 1);
__pyx_t_19 = 0;
__pyx_t_21 = (__pyx_v_c + (__pyx_v_kc - __pyx_v_center_c));
__pyx_t_23 = (__pyx_t_19 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_19;
} else {
__pyx_t_20 = __pyx_t_21;
}
__pyx_t_19 = __pyx_t_20;
__pyx_t_23 = (__pyx_t_22 < __pyx_t_19);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_22;
} else {
__pyx_t_20 = __pyx_t_19;
}
__pyx_v_local_col = __pyx_t_20;
+168: acc = acc + kernel[kr,kc] * image[local_row, local_col]
__pyx_t_24 = __pyx_v_kr;
__pyx_t_25 = __pyx_v_kc;
__pyx_t_26 = __pyx_v_local_row;
__pyx_t_27 = __pyx_v_local_col;
__pyx_v_acc = (__pyx_v_acc + ((*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_kernel.data + __pyx_t_24 * __pyx_v_kernel.strides[0]) ) + __pyx_t_25 * __pyx_v_kernel.strides[1]) ))) * (*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_26 * __pyx_v_image.strides[0]) ) + __pyx_t_27 * __pyx_v_image.strides[1]) )))));
}
}
+169: _conv_out[r,c] = acc
__pyx_t_27 = __pyx_v_r;
__pyx_t_26 = __pyx_v_c;
*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v__conv_out.data + __pyx_t_27 * __pyx_v__conv_out.strides[0]) ) + __pyx_t_26 * __pyx_v__conv_out.strides[1]) )) = __pyx_v_acc;
}
}
}
}
}
}
}
}
}
}
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) __builtin_expect(!!(x), 1)
#define unlikely(x) __builtin_expect(!!(x), 0)
#endif
}
170:
+171: return conv_out
__Pyx_XDECREF(__pyx_r); __Pyx_INCREF(__pyx_v_conv_out); __pyx_r = __pyx_v_conv_out; goto __pyx_L0;
172:
+173: def _run_threaded_guided(self, float[:,:] image, float[:,:] kernel):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_15_run_threaded_guided(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_15_run_threaded_guided = {"_run_threaded_guided", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_15_run_threaded_guided, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_15_run_threaded_guided(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_kernel = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_guided (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 173, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_kernel,0};
PyObject* values[3] = {0,0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 173, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 173, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded_guided", 1, 3, 3, 1); __PYX_ERR(0, 173, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (likely((values[2] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_kernel)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[2]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 173, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded_guided", 1, 3, 3, 2); __PYX_ERR(0, 173, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_threaded_guided") < 0)) __PYX_ERR(0, 173, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 3)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 173, __pyx_L3_error)
__pyx_v_kernel = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(values[2], PyBUF_WRITABLE); if (unlikely(!__pyx_v_kernel.memview)) __PYX_ERR(0, 173, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_threaded_guided", 1, 3, 3, __pyx_nargs); __PYX_ERR(0, 173, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_threaded_guided", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_14_run_threaded_guided(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_kernel);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_kernel, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_14_run_threaded_guided(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image, __Pyx_memviewslice __pyx_v_kernel) {
int __pyx_v_nRows;
int __pyx_v_nCols;
int __pyx_v_nRows_kernel;
int __pyx_v_nCols_kernel;
int __pyx_v_center_r;
int __pyx_v_center_c;
int __pyx_v_r;
int __pyx_v_c;
int __pyx_v_kr;
int __pyx_v_kc;
int __pyx_v_local_row;
int __pyx_v_local_col;
float __pyx_v_acc;
PyObject *__pyx_v_conv_out = NULL;
__Pyx_memviewslice __pyx_v__conv_out = { 0, 0, { 0 }, { 0 }, { 0 } };
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_guided", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__Pyx_XDECREF(__pyx_t_5);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_threaded_guided", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XDECREF(__pyx_v_conv_out);
__PYX_XCLEAR_MEMVIEW(&__pyx_v__conv_out, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_15_run_threaded_guided, 0, __pyx_n_s_Convolution__run_threaded_guided, NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_d, ((PyObject *)__pyx_codeobj__36)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 173, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_threaded_guided, __pyx_t_10) < 0) __PYX_ERR(0, 173, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
/* … */
__pyx_codeobj__36 = (PyObject*)__Pyx_PyCode_New(3, 0, 0, 18, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__31, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_convoluti, __pyx_n_s_run_threaded_guided, 173, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__36)) __PYX_ERR(0, 173, __pyx_L1_error)
174:
+175: cdef int nRows = image.shape[0]
__pyx_v_nRows = (__pyx_v_image.shape[0]);
+176: cdef int nCols = image.shape[1]
__pyx_v_nCols = (__pyx_v_image.shape[1]);
177:
+178: cdef int nRows_kernel = kernel.shape[0]
__pyx_v_nRows_kernel = (__pyx_v_kernel.shape[0]);
+179: cdef int nCols_kernel = kernel.shape[1]
__pyx_v_nCols_kernel = (__pyx_v_kernel.shape[1]);
180:
+181: cdef int center_r = (nRows_kernel-1) // 2
__pyx_v_center_r = ((__pyx_v_nRows_kernel - 1) / 2);
+182: cdef int center_c = (nCols_kernel-1) // 2
__pyx_v_center_c = ((__pyx_v_nCols_kernel - 1) / 2);
183:
184: cdef int r,c
185: cdef int kr,kc
186:
187: cdef int local_row, local_col
188:
+189: cdef float acc = 0
__pyx_v_acc = 0.0;
190:
+191: conv_out = np.zeros((nRows, nCols), dtype=np.float32)
__Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_zeros); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyInt_From_int(__pyx_v_nRows); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_3 = __Pyx_PyInt_From_int(__pyx_v_nCols); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = PyTuple_New(2); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 0, __pyx_t_1)) __PYX_ERR(0, 191, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_3); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 1, __pyx_t_3)) __PYX_ERR(0, 191, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_3 = 0; __pyx_t_3 = PyTuple_New(1); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_4); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_4)) __PYX_ERR(0, 191, __pyx_L1_error); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_5 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_float32); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; if (PyDict_SetItem(__pyx_t_4, __pyx_n_s_dtype, __pyx_t_5) < 0) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_DECREF(__pyx_t_5); __pyx_t_5 = 0; __pyx_t_5 = __Pyx_PyObject_Call(__pyx_t_2, __pyx_t_3, __pyx_t_4); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 191, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_v_conv_out = __pyx_t_5; __pyx_t_5 = 0;
+192: cdef float[:,:] _conv_out = conv_out
__pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsds_float(__pyx_v_conv_out, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 192, __pyx_L1_error) __pyx_v__conv_out = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
193:
+194: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+195: for r in prange(nRows, schedule="guided"):
__pyx_t_7 = __pyx_v_nRows;
{
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) (x)
#define unlikely(x) (x)
#endif
__pyx_t_9 = (__pyx_t_7 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_9 > 0)
{
#ifdef _OPENMP
#pragma omp parallel
#endif /* _OPENMP */
{
#ifdef _OPENMP
#pragma omp for lastprivate(__pyx_v_acc) lastprivate(__pyx_v_c) lastprivate(__pyx_v_kc) lastprivate(__pyx_v_kr) lastprivate(__pyx_v_local_col) lastprivate(__pyx_v_local_row) firstprivate(__pyx_v_r) lastprivate(__pyx_v_r) schedule(guided)
#endif /* _OPENMP */
for (__pyx_t_8 = 0; __pyx_t_8 < __pyx_t_9; __pyx_t_8++){
{
__pyx_v_r = (int)(0 + 1 * __pyx_t_8);
/* Initialize private variables to invalid values */
__pyx_v_acc = ((float)__PYX_NAN());
__pyx_v_c = ((int)0xbad0bad0);
__pyx_v_kc = ((int)0xbad0bad0);
__pyx_v_kr = ((int)0xbad0bad0);
__pyx_v_local_col = ((int)0xbad0bad0);
__pyx_v_local_row = ((int)0xbad0bad0);
+196: for c in prange(nCols, schedule="guided"):
__pyx_t_10 = __pyx_v_nCols;
{
__pyx_t_12 = (__pyx_t_10 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_12 > 0)
{
#if 0
#pragma omp parallel
#endif /* _OPENMP */
{
#if 0
#pragma omp for lastprivate(__pyx_v_acc) firstprivate(__pyx_v_c) lastprivate(__pyx_v_c) lastprivate(__pyx_v_kc) lastprivate(__pyx_v_kr) lastprivate(__pyx_v_local_col) lastprivate(__pyx_v_local_row) schedule(guided)
#endif /* _OPENMP */
for (__pyx_t_11 = 0; __pyx_t_11 < __pyx_t_12; __pyx_t_11++){
{
__pyx_v_c = (int)(0 + 1 * __pyx_t_11);
/* Initialize private variables to invalid values */
__pyx_v_acc = ((float)__PYX_NAN());
__pyx_v_kc = ((int)0xbad0bad0);
__pyx_v_kr = ((int)0xbad0bad0);
__pyx_v_local_col = ((int)0xbad0bad0);
__pyx_v_local_row = ((int)0xbad0bad0);
+197: acc = 0
__pyx_v_acc = 0.0;
+198: for kr in range(nRows_kernel):
__pyx_t_13 = __pyx_v_nRows_kernel;
__pyx_t_14 = __pyx_t_13;
for (__pyx_t_15 = 0; __pyx_t_15 < __pyx_t_14; __pyx_t_15+=1) {
__pyx_v_kr = __pyx_t_15;
+199: for kc in range(nCols_kernel):
__pyx_t_16 = __pyx_v_nCols_kernel;
__pyx_t_17 = __pyx_t_16;
for (__pyx_t_18 = 0; __pyx_t_18 < __pyx_t_17; __pyx_t_18+=1) {
__pyx_v_kc = __pyx_t_18;
+200: local_row = min(max(r+(kr-center_r),0),nRows-1)
__pyx_t_19 = (__pyx_v_nRows - 1);
__pyx_t_20 = 0;
__pyx_t_21 = (__pyx_v_r + (__pyx_v_kr - __pyx_v_center_r));
__pyx_t_23 = (__pyx_t_20 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_20;
} else {
__pyx_t_22 = __pyx_t_21;
}
__pyx_t_20 = __pyx_t_22;
__pyx_t_23 = (__pyx_t_19 < __pyx_t_20);
if (__pyx_t_23) {
__pyx_t_22 = __pyx_t_19;
} else {
__pyx_t_22 = __pyx_t_20;
}
__pyx_v_local_row = __pyx_t_22;
+201: local_col = min(max(c+(kc-center_c),0),nCols-1)
__pyx_t_22 = (__pyx_v_nCols - 1);
__pyx_t_19 = 0;
__pyx_t_21 = (__pyx_v_c + (__pyx_v_kc - __pyx_v_center_c));
__pyx_t_23 = (__pyx_t_19 > __pyx_t_21);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_19;
} else {
__pyx_t_20 = __pyx_t_21;
}
__pyx_t_19 = __pyx_t_20;
__pyx_t_23 = (__pyx_t_22 < __pyx_t_19);
if (__pyx_t_23) {
__pyx_t_20 = __pyx_t_22;
} else {
__pyx_t_20 = __pyx_t_19;
}
__pyx_v_local_col = __pyx_t_20;
+202: acc = acc + kernel[kr,kc] * image[local_row, local_col]
__pyx_t_24 = __pyx_v_kr;
__pyx_t_25 = __pyx_v_kc;
__pyx_t_26 = __pyx_v_local_row;
__pyx_t_27 = __pyx_v_local_col;
__pyx_v_acc = (__pyx_v_acc + ((*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_kernel.data + __pyx_t_24 * __pyx_v_kernel.strides[0]) ) + __pyx_t_25 * __pyx_v_kernel.strides[1]) ))) * (*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_26 * __pyx_v_image.strides[0]) ) + __pyx_t_27 * __pyx_v_image.strides[1]) )))));
}
}
+203: _conv_out[r,c] = acc
__pyx_t_27 = __pyx_v_r;
__pyx_t_26 = __pyx_v_c;
*((float *) ( /* dim=1 */ (( /* dim=0 */ (__pyx_v__conv_out.data + __pyx_t_27 * __pyx_v__conv_out.strides[0]) ) + __pyx_t_26 * __pyx_v__conv_out.strides[1]) )) = __pyx_v_acc;
}
}
}
}
}
}
}
}
}
}
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) __builtin_expect(!!(x), 1)
#define unlikely(x) __builtin_expect(!!(x), 0)
#endif
}
204:
+205: return conv_out
__Pyx_XDECREF(__pyx_r); __Pyx_INCREF(__pyx_v_conv_out); __pyx_r = __pyx_v_conv_out; goto __pyx_L0;
206:
207:
+208: def _run_opencl(self, image, kernel, device):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_17_run_opencl(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_17_run_opencl = {"_run_opencl", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_17_run_opencl, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_15_le_convolution_11Convolution_17_run_opencl(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
PyObject *__pyx_v_self = 0;
PyObject *__pyx_v_image = 0;
PyObject *__pyx_v_kernel = 0;
PyObject *__pyx_v_device = 0;
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_opencl (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 208, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_kernel,&__pyx_n_s_device,0};
PyObject* values[4] = {0,0,0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 4: values[3] = __Pyx_Arg_FASTCALL(__pyx_args, 3);
CYTHON_FALLTHROUGH;
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 208, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 208, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_opencl", 1, 4, 4, 1); __PYX_ERR(0, 208, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (likely((values[2] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_kernel)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[2]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 208, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_opencl", 1, 4, 4, 2); __PYX_ERR(0, 208, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 3:
if (likely((values[3] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_device)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[3]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 208, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_opencl", 1, 4, 4, 3); __PYX_ERR(0, 208, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_opencl") < 0)) __PYX_ERR(0, 208, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 4)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
values[3] = __Pyx_Arg_FASTCALL(__pyx_args, 3);
}
__pyx_v_self = values[0];
__pyx_v_image = values[1];
__pyx_v_kernel = values[2];
__pyx_v_device = values[3];
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_opencl", 1, 4, 4, __pyx_nargs); __PYX_ERR(0, 208, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_opencl", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_16_run_opencl(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_kernel, __pyx_v_device);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_15_le_convolution_11Convolution_16_run_opencl(CYTHON_UNUSED PyObject *__pyx_self, PyObject *__pyx_v_self, PyObject *__pyx_v_image, PyObject *__pyx_v_kernel, PyObject *__pyx_v_device) {
PyObject *__pyx_v_cl_ctx = NULL;
CYTHON_UNUSED PyObject *__pyx_v_dc = NULL;
PyObject *__pyx_v_cl_queue = NULL;
PyObject *__pyx_v_image_out = NULL;
CYTHON_UNUSED PyObject *__pyx_v_mf = NULL;
PyObject *__pyx_v_input_image = NULL;
PyObject *__pyx_v_input_kernel = NULL;
PyObject *__pyx_v_output_opencl = NULL;
PyObject *__pyx_v_code = NULL;
PyObject *__pyx_v_prg = NULL;
PyObject *__pyx_v_knl = NULL;
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_opencl", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__Pyx_XDECREF(__pyx_t_6);
__Pyx_XDECREF(__pyx_t_7);
__Pyx_XDECREF(__pyx_t_8);
__Pyx_XDECREF(__pyx_t_9);
__Pyx_XDECREF(__pyx_t_10);
__Pyx_AddTraceback("nanopyx.liquid._le_convolution.Convolution._run_opencl", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XDECREF(__pyx_v_cl_ctx);
__Pyx_XDECREF(__pyx_v_dc);
__Pyx_XDECREF(__pyx_v_cl_queue);
__Pyx_XDECREF(__pyx_v_image_out);
__Pyx_XDECREF(__pyx_v_mf);
__Pyx_XDECREF(__pyx_v_input_image);
__Pyx_XDECREF(__pyx_v_input_kernel);
__Pyx_XDECREF(__pyx_v_output_opencl);
__Pyx_XDECREF(__pyx_v_code);
__Pyx_XDECREF(__pyx_v_prg);
__Pyx_XDECREF(__pyx_v_knl);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__37 = PyTuple_Pack(15, __pyx_n_s_self, __pyx_n_s_image, __pyx_n_s_kernel, __pyx_n_s_device, __pyx_n_s_cl_ctx, __pyx_n_s_dc, __pyx_n_s_cl_queue, __pyx_n_s_image_out, __pyx_n_s_mf, __pyx_n_s_input_image, __pyx_n_s_input_kernel, __pyx_n_s_output_opencl, __pyx_n_s_code, __pyx_n_s_prg, __pyx_n_s_knl); if (unlikely(!__pyx_tuple__37)) __PYX_ERR(0, 208, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__37);
__Pyx_GIVEREF(__pyx_tuple__37);
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_15_le_convolution_11Convolution_17_run_opencl, 0, __pyx_n_s_Convolution__run_opencl, NULL, __pyx_n_s_nanopyx_liquid__le_convolution, __pyx_d, ((PyObject *)__pyx_codeobj__38)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 208, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_opencl, __pyx_t_10) < 0) __PYX_ERR(0, 208, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
209:
210: # QUEUE AND CONTEXT
+211: cl_ctx = cl.Context([device['device']])
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_cl); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 211, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_Context); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 211, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = __Pyx_PyObject_Dict_GetItem(__pyx_v_device, __pyx_n_u_device); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 211, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_4 = PyList_New(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 211, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyList_SET_ITEM(__pyx_t_4, 0, __pyx_t_2)) __PYX_ERR(0, 211, __pyx_L1_error); __pyx_t_2 = 0; __pyx_t_2 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_3))) { __pyx_t_2 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_2)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_2); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_2, __pyx_t_4}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 211, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_v_cl_ctx = __pyx_t_1; __pyx_t_1 = 0;
+212: dc = device['device']
__pyx_t_1 = __Pyx_PyObject_Dict_GetItem(__pyx_v_device, __pyx_n_u_device); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 212, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_v_dc = __pyx_t_1; __pyx_t_1 = 0;
+213: cl_queue = cl.CommandQueue(cl_ctx)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_cl); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 213, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_CommandQueue); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 213, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_4))) { __pyx_t_3 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_3)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_3); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_3, __pyx_v_cl_ctx}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 213, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __pyx_v_cl_queue = __pyx_t_1; __pyx_t_1 = 0;
214:
+215: image_out = np.zeros((image.shape[0], image.shape[1]), dtype=np.float32)
__Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_np); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_zeros); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_v_image, __pyx_n_s_shape); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_3 = __Pyx_GetItemInt(__pyx_t_1, 0, long, 1, __Pyx_PyInt_From_long, 0, 0, 0); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_v_image, __pyx_n_s_shape); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __Pyx_GetItemInt(__pyx_t_1, 1, long, 1, __Pyx_PyInt_From_long, 0, 0, 0); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = PyTuple_New(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_GIVEREF(__pyx_t_3); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_1, 0, __pyx_t_3)) __PYX_ERR(0, 215, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_1, 1, __pyx_t_2)) __PYX_ERR(0, 215, __pyx_L1_error); __pyx_t_3 = 0; __pyx_t_2 = 0; __pyx_t_2 = PyTuple_New(1); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_2, 0, __pyx_t_1)) __PYX_ERR(0, 215, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_np); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_6 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_float32); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_dtype, __pyx_t_6) < 0) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; __pyx_t_6 = __Pyx_PyObject_Call(__pyx_t_4, __pyx_t_2, __pyx_t_1); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 215, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_image_out = __pyx_t_6; __pyx_t_6 = 0;
+216: mf = cl.mem_flags
__Pyx_GetModuleGlobalName(__pyx_t_6, __pyx_n_s_cl); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 216, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_t_6, __pyx_n_s_mem_flags); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 216, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; __pyx_v_mf = __pyx_t_1; __pyx_t_1 = 0;
217:
+218: input_image = cl.image_from_array(cl_ctx, image, mode='r')
__Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_cl); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 218, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_6 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_image_from_array); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 218, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = PyTuple_New(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 218, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_INCREF(__pyx_v_cl_ctx); __Pyx_GIVEREF(__pyx_v_cl_ctx); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_1, 0, __pyx_v_cl_ctx)) __PYX_ERR(0, 218, __pyx_L1_error); __Pyx_INCREF(__pyx_v_image); __Pyx_GIVEREF(__pyx_v_image); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_1, 1, __pyx_v_image)) __PYX_ERR(0, 218, __pyx_L1_error); __pyx_t_2 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 218, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); if (PyDict_SetItem(__pyx_t_2, __pyx_n_s_mode, __pyx_n_u_r) < 0) __PYX_ERR(0, 218, __pyx_L1_error) __pyx_t_4 = __Pyx_PyObject_Call(__pyx_t_6, __pyx_t_1, __pyx_t_2); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 218, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_v_input_image = __pyx_t_4; __pyx_t_4 = 0;
+219: input_kernel = cl.image_from_array(cl_ctx, kernel, mode='r')
__Pyx_GetModuleGlobalName(__pyx_t_4, __pyx_n_s_cl); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 219, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_4, __pyx_n_s_image_from_array); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 219, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = PyTuple_New(2); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 219, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_INCREF(__pyx_v_cl_ctx); __Pyx_GIVEREF(__pyx_v_cl_ctx); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 0, __pyx_v_cl_ctx)) __PYX_ERR(0, 219, __pyx_L1_error); __Pyx_INCREF(__pyx_v_kernel); __Pyx_GIVEREF(__pyx_v_kernel); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 1, __pyx_v_kernel)) __PYX_ERR(0, 219, __pyx_L1_error); __pyx_t_1 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 219, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_mode, __pyx_n_u_r) < 0) __PYX_ERR(0, 219, __pyx_L1_error) __pyx_t_6 = __Pyx_PyObject_Call(__pyx_t_2, __pyx_t_4, __pyx_t_1); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 219, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_input_kernel = __pyx_t_6; __pyx_t_6 = 0;
+220: output_opencl = cl.image_from_array(cl_ctx, image_out, mode='w')
__Pyx_GetModuleGlobalName(__pyx_t_6, __pyx_n_s_cl); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 220, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_t_6, __pyx_n_s_image_from_array); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 220, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; __pyx_t_6 = PyTuple_New(2); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 220, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __Pyx_INCREF(__pyx_v_cl_ctx); __Pyx_GIVEREF(__pyx_v_cl_ctx); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_6, 0, __pyx_v_cl_ctx)) __PYX_ERR(0, 220, __pyx_L1_error); __Pyx_INCREF(__pyx_v_image_out); __Pyx_GIVEREF(__pyx_v_image_out); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_6, 1, __pyx_v_image_out)) __PYX_ERR(0, 220, __pyx_L1_error); __pyx_t_4 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 220, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); if (PyDict_SetItem(__pyx_t_4, __pyx_n_s_mode, __pyx_n_u_w) < 0) __PYX_ERR(0, 220, __pyx_L1_error) __pyx_t_2 = __Pyx_PyObject_Call(__pyx_t_1, __pyx_t_6, __pyx_t_4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 220, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_v_output_opencl = __pyx_t_2; __pyx_t_2 = 0;
+221: cl_queue.finish()
__pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_cl_queue, __pyx_n_s_finish); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 221, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_6 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_4))) { __pyx_t_6 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_6)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_6); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_6, }; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_6); __pyx_t_6 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 221, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0;
222:
+223: code = self._get_cl_code("_le_convolution.cl", device['DP'])
__pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_self, __pyx_n_s_get_cl_code); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 223, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_6 = __Pyx_PyObject_Dict_GetItem(__pyx_v_device, __pyx_n_u_DP); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 223, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __pyx_t_1 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_4))) { __pyx_t_1 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_1)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_1); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[3] = {__pyx_t_1, __pyx_kp_u_le_convolution_cl, __pyx_t_6}; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 2+__pyx_t_5); __Pyx_XDECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 223, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __pyx_v_code = __pyx_t_2; __pyx_t_2 = 0;
+224: prg = cl.Program(cl_ctx, code).build()
__Pyx_GetModuleGlobalName(__pyx_t_6, __pyx_n_s_cl); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 224, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_t_6, __pyx_n_s_Program); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 224, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; __pyx_t_6 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_1))) { __pyx_t_6 = PyMethod_GET_SELF(__pyx_t_1); if (likely(__pyx_t_6)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_1); __Pyx_INCREF(__pyx_t_6); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_1, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[3] = {__pyx_t_6, __pyx_v_cl_ctx, __pyx_v_code}; __pyx_t_4 = __Pyx_PyObject_FastCall(__pyx_t_1, __pyx_callargs+1-__pyx_t_5, 2+__pyx_t_5); __Pyx_XDECREF(__pyx_t_6); __pyx_t_6 = 0; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 224, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; } __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_t_4, __pyx_n_s_build); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 224, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_1))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_1); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_1); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_1, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_4, }; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_1, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 224, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; } __pyx_v_prg = __pyx_t_2; __pyx_t_2 = 0;
+225: knl = prg.conv2d
__pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_v_prg, __pyx_n_s_conv2d); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 225, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_v_knl = __pyx_t_2; __pyx_t_2 = 0;
226:
227: knl(cl_queue,
+228: (1,image.shape[0], image.shape[1]),
__pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_image, __pyx_n_s_shape); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 228, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_6 = __Pyx_GetItemInt(__pyx_t_4, 0, long, 1, __Pyx_PyInt_From_long, 0, 0, 0); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 228, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_image, __pyx_n_s_shape); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 228, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_3 = __Pyx_GetItemInt(__pyx_t_4, 1, long, 1, __Pyx_PyInt_From_long, 0, 0, 0); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 228, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = PyTuple_New(3); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 228, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_INCREF(__pyx_int_1); __Pyx_GIVEREF(__pyx_int_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 0, __pyx_int_1)) __PYX_ERR(0, 228, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_6); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 1, __pyx_t_6)) __PYX_ERR(0, 228, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_3); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 2, __pyx_t_3)) __PYX_ERR(0, 228, __pyx_L1_error); __pyx_t_6 = 0; __pyx_t_3 = 0;
+229: self.get_work_group(device['device'],(1,image.shape[0], image.shape[1])),
__pyx_t_6 = __Pyx_PyObject_GetAttrStr(__pyx_v_self, __pyx_n_s_get_work_group); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 229, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __pyx_t_7 = __Pyx_PyObject_Dict_GetItem(__pyx_v_device, __pyx_n_u_device); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 229, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __pyx_t_8 = __Pyx_PyObject_GetAttrStr(__pyx_v_image, __pyx_n_s_shape); if (unlikely(!__pyx_t_8)) __PYX_ERR(0, 229, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_8); __pyx_t_9 = __Pyx_GetItemInt(__pyx_t_8, 0, long, 1, __Pyx_PyInt_From_long, 0, 0, 0); if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 229, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); __Pyx_DECREF(__pyx_t_8); __pyx_t_8 = 0; __pyx_t_8 = __Pyx_PyObject_GetAttrStr(__pyx_v_image, __pyx_n_s_shape); if (unlikely(!__pyx_t_8)) __PYX_ERR(0, 229, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_8); __pyx_t_10 = __Pyx_GetItemInt(__pyx_t_8, 1, long, 1, __Pyx_PyInt_From_long, 0, 0, 0); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 229, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_10); __Pyx_DECREF(__pyx_t_8); __pyx_t_8 = 0; __pyx_t_8 = PyTuple_New(3); if (unlikely(!__pyx_t_8)) __PYX_ERR(0, 229, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_8); __Pyx_INCREF(__pyx_int_1); __Pyx_GIVEREF(__pyx_int_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_8, 0, __pyx_int_1)) __PYX_ERR(0, 229, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_9); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_8, 1, __pyx_t_9)) __PYX_ERR(0, 229, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_10); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_8, 2, __pyx_t_10)) __PYX_ERR(0, 229, __pyx_L1_error); __pyx_t_9 = 0; __pyx_t_10 = 0; __pyx_t_10 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_6))) { __pyx_t_10 = PyMethod_GET_SELF(__pyx_t_6); if (likely(__pyx_t_10)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_6); __Pyx_INCREF(__pyx_t_10); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_6, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[3] = {__pyx_t_10, __pyx_t_7, __pyx_t_8}; __pyx_t_3 = __Pyx_PyObject_FastCall(__pyx_t_6, __pyx_callargs+1-__pyx_t_5, 2+__pyx_t_5); __Pyx_XDECREF(__pyx_t_10); __pyx_t_10 = 0; __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __Pyx_DECREF(__pyx_t_8); __pyx_t_8 = 0; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 229, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; }
230: input_image,
231: output_opencl,
+232: input_kernel).wait()
__Pyx_INCREF(__pyx_v_knl); __pyx_t_6 = __pyx_v_knl; __pyx_t_8 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_6))) { __pyx_t_8 = PyMethod_GET_SELF(__pyx_t_6); if (likely(__pyx_t_8)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_6); __Pyx_INCREF(__pyx_t_8); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_6, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[7] = {__pyx_t_8, __pyx_v_cl_queue, __pyx_t_4, __pyx_t_3, __pyx_v_input_image, __pyx_v_output_opencl, __pyx_v_input_kernel}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_6, __pyx_callargs+1-__pyx_t_5, 6+__pyx_t_5); __Pyx_XDECREF(__pyx_t_8); __pyx_t_8 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 227, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; } __pyx_t_6 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_wait); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 232, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_6))) { __pyx_t_1 = PyMethod_GET_SELF(__pyx_t_6); if (likely(__pyx_t_1)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_6); __Pyx_INCREF(__pyx_t_1); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_6, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_1, }; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_6, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_1); __pyx_t_1 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 232, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; } __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0;
233:
+234: cl.enqueue_copy(cl_queue, image_out, output_opencl,origin=(0,0), region=(image.shape[0], image.shape[1])).wait()
__Pyx_GetModuleGlobalName(__pyx_t_6, __pyx_n_s_cl); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_t_6, __pyx_n_s_enqueue_copy); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; __pyx_t_6 = PyTuple_New(3); if (unlikely(!__pyx_t_6)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_6); __Pyx_INCREF(__pyx_v_cl_queue); __Pyx_GIVEREF(__pyx_v_cl_queue); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_6, 0, __pyx_v_cl_queue)) __PYX_ERR(0, 234, __pyx_L1_error); __Pyx_INCREF(__pyx_v_image_out); __Pyx_GIVEREF(__pyx_v_image_out); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_6, 1, __pyx_v_image_out)) __PYX_ERR(0, 234, __pyx_L1_error); __Pyx_INCREF(__pyx_v_output_opencl); __Pyx_GIVEREF(__pyx_v_output_opencl); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_6, 2, __pyx_v_output_opencl)) __PYX_ERR(0, 234, __pyx_L1_error); __pyx_t_3 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); if (PyDict_SetItem(__pyx_t_3, __pyx_n_s_origin, __pyx_tuple__11) < 0) __PYX_ERR(0, 234, __pyx_L1_error) __pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_image, __pyx_n_s_shape); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_8 = __Pyx_GetItemInt(__pyx_t_4, 0, long, 1, __Pyx_PyInt_From_long, 0, 0, 0); if (unlikely(!__pyx_t_8)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_8); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_image, __pyx_n_s_shape); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_7 = __Pyx_GetItemInt(__pyx_t_4, 1, long, 1, __Pyx_PyInt_From_long, 0, 0, 0); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = PyTuple_New(2); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GIVEREF(__pyx_t_8); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 0, __pyx_t_8)) __PYX_ERR(0, 234, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_7); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_4, 1, __pyx_t_7)) __PYX_ERR(0, 234, __pyx_L1_error); __pyx_t_8 = 0; __pyx_t_7 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_s_region, __pyx_t_4) < 0) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PyObject_Call(__pyx_t_1, __pyx_t_6, __pyx_t_3); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_6); __pyx_t_6 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_4, __pyx_n_s_wait); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_3))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_4, }; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; /* … */ __pyx_tuple__11 = PyTuple_Pack(2, __pyx_int_0, __pyx_int_0); if (unlikely(!__pyx_tuple__11)) __PYX_ERR(0, 234, __pyx_L1_error) __Pyx_GOTREF(__pyx_tuple__11); __Pyx_GIVEREF(__pyx_tuple__11);
235:
+236: cl_queue.finish()
__pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_v_cl_queue, __pyx_n_s_finish); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 236, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_3))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_4, }; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 236, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0;
+237: return image_out
__Pyx_XDECREF(__pyx_r); __Pyx_INCREF(__pyx_v_image_out); __pyx_r = __pyx_v_image_out; goto __pyx_L0;